You write custom CUDA kernels to replace pytorch operators in given architecture to get speedups. You have complete freedom to choose the set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into a single kernel, for example, combining matmul+relu), or algorithmic changes (such as online softmax). You are only limited by your imagination.

Here's an example to show you the syntax of inline embedding custom CUDA operators in PyTorch. The example given architecture is a simple addition:

python
import torch
import torch.nn as nn
import torch.nn.functional as F

class Model(nn.Module):
def init(self) -> None:
super().init()

def forward(self, a, b):
    return a + b
def get_inputs():
a = torch.randn(1, 128).cuda()
b = torch.randn(1, 128).cuda()
return [a, b]

def get_init_inputs():
return []



The example new architecture with a custom CUDA kernel looks like this:

python
import torch
from torch.utils.cpp_extension import load_inline

add_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

global void add_kernel(const float* a, const float* b, float* out, int size) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
out[idx] = a[idx] + b[idx];
}
}

torch::Tensor add_cuda(torch::Tensor a, torch::Tensor b) {
auto out = torch::empty_like(a);
int size = a.numel();
const int block_size = 256;
int num_blocks = (size + block_size - 1) / block_size;
add_kernel<<<num_blocks, block_size>>>(a.data_ptr<float>(), b.data_ptr<float>(), out.data_ptr<float>(), size);
return out;
}
"""

add_cpp_source = """
torch::Tensor add_cuda(torch::Tensor a, torch::Tensor b);
"""

Compile the inline CUDA code
add = load_inline(
name="add",
cpp_sources=add_cpp_source,
cuda_sources=add_source,
functions=["add_cuda"],
verbose=True
)

class ModelNew(torch.nn.Module):
def init(self):
super(ModelNew, self).init()
self.add = add

def forward(self, a, b):
    return self.add.add_cuda(a, b)


---

Now, you are given the following PyTorch architecture to accelerate. The model computes the Manhattan distance (L1 distance) between two sets of vectors and then applies the Hard-Swish activation function to each distance. This baseline implementation is efficient and uses PyTorch's highly optimized built-in functions for correctness and performance.

python
import torch
import torch.nn as nn

class Model(nn.Module):
    """
    PyTorch基准实现：曼哈顿距离 + Hard-Swish激活
    """
    def __init__(self):
        super(Model, self).__init__()
    
    def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
        """
        Compute Manhattan distance between x and y, then apply Hard-Swish activation.

        Args:
            x (torch.Tensor): First set of vectors [batch_size, feature_dim]
            y (torch.Tensor): Second set of vectors [batch_size, feature_dim]

        Returns:
            torch.Tensor: Hard-Swish-activated Manhattan distances [batch_size]
        """
        # Input validation
        if x.shape != y.shape:
            raise ValueError(f"Input tensors must have the same shape, got {x.shape} and {y.shape}")
        
        if x.dim() != 2:
            raise ValueError(f"Input tensors must be 2D, got {x.dim()}D")
        
        # --- 第一步：计算曼哈顿距离 ---
        # Σ|x_i - y_i|
        manhattan_dist = torch.sum(torch.abs(x - y), dim=1)
        
        # --- 第二步：应用Hard-Swish激活函数 ---
        # 使用F.hardswish，这是PyTorch内置的高效实现
        activated_distances = torch.nn.functional.hardswish(manhattan_dist)
        
        return activated_distances

batch_size = 256
feature_dim = 512

def get_inputs():
    # Generate two sets of vectors
    x = torch.randn(batch_size, feature_dim)
    y = torch.randn(batch_size, feature_dim)
    return [x, y]

def get_init_inputs():
    return []  # No special initialization inputs needed



Your task is to generate the `ModelNew` architecture with a custom CUDA kernel that computes the Manhattan distance and applies the Hard-Swish activation in a fused manner. The implementation must be highly optimized.

**CRITICAL REQUIREMENTS:**

1.  **Performance Optimization & Fusion:**
    *   **Operator Fusion:** The entire calculation (computing the Manhattan distance and then applying the Hard-Swish activation function for each sample) must be performed within a **single CUDA kernel**. This kernel should output a tensor of Hard-Swish-activated distances.
    *   The kernel should use a **multi-threaded reduction with float4 vectorization** strategy within each block to sum the absolute differences. After the reduction, a single thread should apply the Hard-Swish activation and store the result.

2.  **Kernel Logic:**
    *   The kernel should launch a 1D grid where each block corresponds to one sample in the batch.
    *   Use `float4` for vectorized memory access to improve bandwidth utilization. Each thread should process 4 elements at a time.
    *   Use shared memory to store the partial sums found by each thread, and then perform a parallel reduction within the block to find the total sum for that sample.
    *   The Hard-Swish activation is a piecewise function: `0` if `x <= -3`, `x` if `x >= 3`, and `x * (x + 3) / 6` otherwise. **Implement this efficiently using the ternary conditional operator (`?:`) to avoid branch divergence.**

3.  **Code Structure:** Follow the exact structure of the provided example, including `load_inline`, the CUDA source string, the C++ wrapper source string, and the `ModelNew` class. The `get_init_inputs` function must return `[]` to match the baseline.

4.  **Compilation Flags:** Use `-O3` for optimization but **do not** use `--use_fast_math` to ensure numerical accuracy with the PyTorch baseline. Avoid hardcoding compute capabilities to ensure portability.
